<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<title>Mali OpenCL SDK v1.1.0: Sobel Filter</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
  $(document).ready(initResizable);
</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td id="projectlogo"><img alt="Logo" src="logo.png"/></td>
  <td style="padding-left: 0.5em;">
   <div id="projectname">Mali OpenCL SDK v1.1.0
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.2 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Home</span></a></li>
      <li class="current"><a href="pages.html"><span>Help&#160;and&#160;Tutorials</span></a></li>
      <li><a href="files.html"><span>Files</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
</div><!-- top -->
<div id="side-nav" class="ui-resizable side-nav-resizable">
  <div id="nav-tree">
    <div id="nav-tree-contents">
      <div id="nav-sync" class="sync"></div>
    </div>
  </div>
  <div id="splitbar" style="-moz-user-select:none;" 
       class="ui-resizable-handle">
  </div>
</div>
<script type="text/javascript">
$(document).ready(function(){initNavTree('sobel_tutorial.html','');});
</script>
<div id="doc-content">
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Macros</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Pages</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div class="header">
  <div class="headertitle">
<div class="title">Sobel Filter </div>  </div>
</div><!--header-->
<div class="contents">
<div class="textblock"><p>The Sobel image filter is a simple convolution filter used primarily for edge detection algorithms.</p>
<h1><a class="anchor" id="sobelExampleResult"></a>
Example Result</h1>
<div class="image">
<img src="sobel_input.bmp" alt="sobel_input.bmp"/>
<div class="caption">
Input image</div></div>
 <div class="image">
<img src="sobel_output.bmp" alt="sobel_output.bmp"/>
<div class="caption">
Output image</div></div>
 <h1><a class="anchor" id="sobelAlgorithm"></a>
The Algorithm</h1>
<p>One technique for doing edge detection on an image is to find the gradient of the image. Areas with large gradients correspond to areas of large change in colour/intensity of the image. These areas are typically edges.</p>
<p>If you <a href="http://en.wikipedia.org/wiki/Convolution">convolve</a> the two Sobel operators with an image, you get two outputs:</p>
<ul>
<li>the gradient in the X direction (dX)</li>
<li>the gradient in the Y direction (dY)</li>
</ul>
<p>The gradients are then typically combined to give a total gradient image.</p>
<p>The Sobel operators are:</p>
<p>dX: </p>
<pre class="fragment"> -1  0  1
 -2  0  2
 -1  0  1
</pre><p> dY: </p>
<pre class="fragment">  1  2  1
  0  0  0
 -1 -2 -1
</pre><p> Convolution is well suited to parallelization. Each output pixel only depends on the constant input pixels, not on outputs from any of the other pixels. Therefore, each pixel can be calculated independently and simultaneously.</p>
<p>For more details see <a href="http://en.wikipedia.org/wiki/Sobel_operator">Wikipedia</a>.</p>
<h1><a class="anchor" id="sobelImplementation"></a>
Implementation</h1>
<h2><a class="anchor" id="sobelSize"></a>
Image Size</h2>
<p>We have included a 512x512 input bitmap for use with this sample (to keep the size of the installer small). However, you are more likely to see performance improvements (when compared to C code running on a CPU) when larger images are used. There is some start-up overhead associated with using OpenCL. This overhead can outweigh the benefits of parallel processing when the input data sizes are small.</p>
<p>This sample has been coded to allow any input bitmap to be used. Simply change <em>input.bmp</em> in the assets directory of the sample to the input image of your choice. You will see larger calculation performance improvements when larger images are used.</p>
<h2><a class="anchor" id="sobelPadding"></a>
Padding</h2>
<p>It is important to note that we have not considered padding here. The output image is two pixels smaller in both dimensions. Because every output requires pixels around it, it is impossible to calculate the output for the edge pixels. In this example we are simply leaving the edge output pixels as the values they are initialized to.</p>
<p>Sometimes it can be desirable to have the size of the output-signal be the same as the size of the input-signal, in which case "padding" must be applied to the input to take into account for the fact that the filter-application, by its nature, reduces the size. Strategies for padding differ, but for images, a common choice is to repeat the boundary-values (i.e. the outmost set of pixels) on all sides or (in some cases) on just some sides.</p>
<h2><a class="anchor" id="sobelSpecifics"></a>
Implementation Specifics</h2>
<p>We implement Sobel filtering on a single 8-bit channel for simplicity. To do Sobel filtering on RGB images you can run the sobel filter on each channel seperately and then combine the results.</p>
<p>In this sample we take an RGB image, convert it to a 8-bit luminance image and send it to the GPU.</p>
<p>Each Sobel calculation gives an output for the centre pixel of the mask. The output value of the centre pixel is the sum of the pixel values in a 3x3 grid around the pixel, multiplied by the Sobel mask. This can be split into three stages by doing the multiplication and summations for each row of the grid separately (the approach used in this sample).</p>
<h2><a class="anchor" id="sobelCode"></a>
The Code</h2>
<dl class="section note"><dt>Note</dt><dd>We consider vectorization the main optimisation technique used in this tutorial. To show the importance of using vectors, an OpenCL implementation of the Sobel filter which doesn't include vectors is included in the sobel_no_vectors sample(<a class="el" href="sobel__no__vectors_8cpp.html">sobel_no_vectors.cpp</a> and <a class="el" href="sobel__no__vectors_8cl.html">sobel_no_vectors.cl</a>). The vectorized version is much faster on Mali-T600 series GPUs.</dd></dl>
<p>Unless otherwise noted, all code snippets come from the OpenCL kernel found in <a class="el" href="sobel_8cl.html">sobel.cl</a>.</p>
<ol type="1">
<li><p class="startli"><b>Choosing the size of the kernel</b></p>
<p class="startli">We are using vector types in the kernel and so we are actually outputting 16 results per kernel. See below for more details of vectorising. The kernel applies the 3x3 Sobel filter to a 18x3 window in the input image to produce 16x1 results in the two ouput images representing the dx and dy components of the gradient. We adjust the pointers into the data to reflect this: </p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Each kernel calculates 16 output pixels in the same row (hence the &#39;* 16&#39;).</span></div>
<div class="line"><span class="comment">     * column is in the range [0, width] in steps of 16.</span></div>
<div class="line"><span class="comment">     * row is in the range [0, height].</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keyword">const</span> <span class="keywordtype">int</span> column = get_global_id(0) * 16;</div>
<div class="line">    <span class="keyword">const</span> <span class="keywordtype">int</span> row = get_global_id(1) * 1;</div>
<div class="line"></div>
<div class="line">    <span class="comment">/* Offset calculates the position in the linear data for the row and the column. */</span></div>
<div class="line">    <span class="keyword">const</span> <span class="keywordtype">int</span> offset = row * width + column;</div>
</div><!-- fragment --><p> And when we enqueue the kernel in <a class="el" href="sobel_8cpp.html">sobel.cpp</a>, we reduce the worksize accordingly: </p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Each instance of the kernel operates on a 16 * 1 portion of the image.</span></div>
<div class="line"><span class="comment">     * Therefore, the global work size must be width / 16 by height / 1 work items.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordtype">size_t</span> globalWorksize[2] = {width / 16, height / 1};</div>
</div><!-- fragment --></li>
<li><p class="startli"><b>Loading the input data</b></p>
<p class="startli">Mali-T600 series GPUs have 128-bit vector registers and can do arithmetic on vector types. Therefore, we use OpenCL vectors to make more efficient use of the hardware, leading to higher performance.</p>
<p class="startli">Here we do vector loads from one row of the data: </p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * First row of input.</span></div>
<div class="line"><span class="comment">     * In a scalar Sobel calculation you would load 1 value for leftLoad, middleLoad and rightLoad.</span></div>
<div class="line"><span class="comment">     * In the vector case we load 16 values for each.</span></div>
<div class="line"><span class="comment">     * leftLoad, middleLoad and rightLoad load 16-bytes of data from the first row.</span></div>
<div class="line"><span class="comment">     * The data they load overlaps. e.g. for the first column and row, leftLoad is 0-&gt;15, middleLoad is 1-&gt;16 and rightLoad is 2-&gt;17.</span></div>
<div class="line"><span class="comment">     * So we&#39;re actually loading 18-bytes of data from the first row.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    uchar16 leftLoad = vload16(0, inputImage + (offset + 0));</div>
<div class="line">    uchar16 middleLoad = vload16(0, inputImage + (offset + 1));</div>
<div class="line">    uchar16 rightLoad = vload16(0, inputImage + (offset + 2));</div>
</div><!-- fragment --></li>
<li><p class="startli"><b>Converting the data</b></p>
<p class="startli">On a Mali-T600 series GPU, expanding and contracting data types is a free operation. Here, we convert the data from 8-bits per pixel to 16-bits per pixel: </p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Convert the data from unsigned chars to shorts (8-bit unsigned to 16-bit signed).</span></div>
<div class="line"><span class="comment">     * The calculations can overflow 8-bits so we require larger intermediate storage.</span></div>
<div class="line"><span class="comment">     * Additionally, the values can become negative so we need a signed type.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    short16 leftData = convert_short16(leftLoad);</div>
<div class="line">    short16 middleData = convert_short16(middleLoad);</div>
<div class="line">    short16 rightData = convert_short16(rightLoad);</div>
</div><!-- fragment --></li>
<li><p class="startli"><b>Doing the calculation</b></p>
<p class="startli">Then we carry out the calculation on 16 pixels. Each vector calculation can be done as a single operation on a Mali-T600 series GPU: </p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Calculate the results for the first row.</span></div>
<div class="line"><span class="comment">     * Looking at the Sobel masks above for the first line of input,</span></div>
<div class="line"><span class="comment">     * the dX calculation is the sum of 1 * leftData, 0 * middleData, and -1 * rightData.</span></div>
<div class="line"><span class="comment">     * The dY calculation is the sum of 1 * leftData, 2 * middleData, and 1 * rightData.</span></div>
<div class="line"><span class="comment">     * This is what is being calculated below, except we have removed the</span></div>
<div class="line"><span class="comment">     * unnecessary calculations (multiplications by 1 or 0) and we are calculating 16 values at once.</span></div>
<div class="line"><span class="comment">     * This pattern repeats for the other 2 rows of data.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    short16 dx = rightData - leftData;</div>
<div class="line">    short16 dy = rightData + leftData + middleData * (short)2;</div>
</div><!-- fragment --><p> We do this for the other two rows, accumulating the results in dx and dy.</p>
</li>
<li><p class="startli"><b>Storing the results</b></p>
<p class="startli">Finally we contract and store the data. We use a vector store to write out all 16 results at once: </p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Store the results.</span></div>
<div class="line"><span class="comment">     * The range of outputs from our Sobel calculations is [-1020, 1020].</span></div>
<div class="line"><span class="comment">     * In order to output this as an 8-bit signed char we must divide it by 8 (or shift right 3 times).</span></div>
<div class="line"><span class="comment">     * This gives the range [-128, 128]. Depending on what type of output you require,</span></div>
<div class="line"><span class="comment">     * (signed/unsigned, seperate/combined gradients) it is possible to do more of the calculations on the GPU using OpenCL.</span></div>
<div class="line"><span class="comment">     * In this sample we&#39;re assuming that the application requires signed uncombined gradient outputs.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    vstore16(convert_char16(dx &gt;&gt; 3), 0, outputImageDX + offset + width + 1);</div>
<div class="line">    vstore16(convert_char16(dy &gt;&gt; 3), 0, outputImageDY + offset + width + 1);</div>
</div><!-- fragment --><p> Because the data is returned as two seperate gradient images we combine them on the CPU before writing them out as a bitmap.</p>
</li>
</ol>
<h1><a class="anchor" id="sobelRunning"></a>
Running the Sample</h1>
<ol type="1">
<li><p class="startli">From a command prompt in the root of the SDK, run:</p>
<div class="fragment"><div class="line">cd samples/<a class="code" href="sobel_8cl.html#ad999d1681236d38b8c3e05c610a99850" title="Sobel filter kernel function.">sobel</a></div>
</div><!-- fragment --> <div class="fragment"><div class="line">make install</div>
</div><!-- fragment --><p class="startli">This compiles the Sobel sample code and copies all the files it needs to run to the bin folder in the root directory of the SDK.</p>
</li>
<li>Copy this folder to the board.</li>
<li><p class="startli">Navigate to the folder on the board and run the Sobel binary:</p>
<div class="fragment"><div class="line">./<a class="code" href="sobel_8cl.html#ad999d1681236d38b8c3e05c610a99850" title="Sobel filter kernel function.">sobel</a></div>
</div><!-- fragment --></li>
<li><p class="startli">You should see output similar to:</p>
<div class="fragment"><div class="line">Profiling information:</div>
<div class="line">Queued time:    0.067ms</div>
<div class="line">Wait time:      0.057835ms</div>
<div class="line">Run time:       0.164923ms</div>
</div><!-- fragment --><p class="startli">An output image should be created on the board called output.bmp.</p>
</li>
</ol>
<p>Find solutions for <a class="el" href="common_issues.html">Common Issues</a>.</p>
<h1><a class="anchor" id="sobelMoreInformation"></a>
More Information</h1>
<p>For more information have a look at the code in <a class="el" href="sobel_8cpp.html">sobel.cpp</a> and <a class="el" href="sobel_8cl.html">sobel.cl</a>. </p>
</div></div><!-- contents -->
</div><!-- doc-content -->
<!-- start footer part -->
<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
  <ul>
    <li class="navelem"><a class="el" href="tutorials.html">Tutorials</a></li>
    <li class="footer">
        <a href="http://www.arm.com/">(C) ARM Ltd. 2013</a>
    </li>
  </ul>
</div>
</body>
</html>
